`:top

`!Parallel Thread Execution`! (`!PTX`! or `!NVPTX`!`:cite-ref-0-1-0[`F5bf`_`[1`#cite-note-0-1]`_`f]) is a low-level `F33f`_`[parallel`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Parallel_computing]`_`f `F33f`_`[thread`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Thread_(computing)]`_`f `F33f`_`[execution`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Execution_(computing)]`_`f `F33f`_`[virtual machine`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Virtual_machine]`_`f and `F33f`_`[instruction set architecture`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Instruction_set_architecture]`_`f used in `F33f`_`[Nvidia`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Nvidia]`_`f's Compute Unified Device Architecture (`F33f`_`[CUDA`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=CUDA]`_`f) programming environment. The `F33f`_`[Nvidia CUDA Compiler`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Nvidia_CUDA_Compiler]`_`f (NVCC) translates code written in CUDA, a `F33f`_`[C++`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=C++]`_`f-like language, into PTX instructions (an `F33f`_`[IL`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Intermediate_language]`_`f), and the graphics driver contains a `F33f`_`[compiler`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Compiler]`_`f which translates PTX instructions into executable binary code,`:cite-ref-2[`F5bf`_`[2`#cite-note-2]`_`f] which can run on the processing cores of `F33f`_`[Nvidia`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=List_of_Nvidia_graphics_processing_units]`_`f `F33f`_`[graphics processing units`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Graphics_processing_unit]`_`f (GPUs). The `F33f`_`[GNU Compiler Collection`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=GNU_Compiler_Collection]`_`f`:cite-ref-3[`F5bf`_`[3`#cite-note-3]`_`f] and `F33f`_`[LLVM`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=LLVM]`_`f`:cite-ref-0-1-1[`F5bf`_`[1`#cite-note-0-1]`_`f] also have the ability to generate PTX. Inline PTX assembly can be used in CUDA.`:cite-ref-4[`F5bf`_`[4`#cite-note-4]`_`f]

>>Contents

• `F0af`_`[Registers`#registers]`_`f

• `F0af`_`[State spaces`#state-spaces]`_`f

• `F0af`_`[See also`#see-also]`_`f

• `F0af`_`[References`#references]`_`f

• `F0af`_`[External links`#external-links]`_`f

-─

>>Registers

PTX uses an arbitrarily large `F33f`_`[processor register`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Processor_register]`_`f set; the output from the compiler is almost pure `F33f`_`[static single-assignment form`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Static_single-assignment_form]`_`f, with consecutive lines generally referring to consecutive registers. Programs start with declarations of the form

`B100`F9d9.reg .u32 %r<335>; // declare 335 registers %r0, %r1, ..., %r334 of type unsigned 32-bit integer`f`b

It is a three-argument assembly language, and almost all instructions explicitly list the data type (in sign and width) on which they operate. Register names are preceded with a % character and constants are literal, e.g.:

`B100`F9d9shr.u64 %rd14, %rd12, 32; // shift right an unsigned 64-bit integer from %rd12 by 32 positions, result in %rd14`f`b

`B100`F9d9cvt.u64.u32 %rd142, %r112; // convert an unsigned 32-bit integer to 64-bit`f`b

There are predicate registers, but compiled code in shader model 1.0 uses these only in conjunction with branch commands; the conditional branch is

`B100`F9d9@%p14 bra $label; // branch to $label`f`b

The `B100`F9d9setp.cc.type`f`b instruction sets a predicate register to the result of comparing two registers of appropriate type, there is also a `B100`F9d9set`f`b instruction, where `B100`F9d9set.le.u32.u64 %r101, %rd12, %rd28`f`b sets the 32-bit register `B100`F9d9%r101`f`b to `B100`F9d90xffffffff`f`b if the 64-bit register `B100`F9d9%rd12`f`b is less than or equal to the 64-bit register `B100`F9d9%rd28`f`b. Otherwise `B100`F9d9%r101`f`b is set to `B100`F9d90x00000000`f`b.

There are a few predefined identifiers that denote pseudoregisters. Among others, `B100`F9d9%tid, %ntid, %ctaid`f`b, and `B100`F9d9%nctaid`f`b contain, respectively, thread indices, block dimensions, block indices, and grid dimensions.`:cite-ref-ptx-isa-5-0[`F5bf`_`[5`#cite-note-ptx-isa-5]`_`f]

>>State spaces

Load (`B100`F9d9ld`f`b) and store (`B100`F9d9st`f`b) commands refer to one of several distinct state spaces (memory banks), e.g. `B100`F9d9ld.param`f`b. There are eight state spaces:`:cite-ref-ptx-isa-5-1[`F5bf`_`[5`#cite-note-ptx-isa-5]`_`f]

`!`B100`F9d9.reg`f`b`!

registers

`!`B100`F9d9.sreg`f`b`!

special, read-only, platform-specific registers

`!`B100`F9d9.const`f`b`!

shared, read-only memory

`!`B100`F9d9.global`f`b`!

global memory, shared by all threads

`!`B100`F9d9.local`f`b`!

local memory, private to each thread

`!`B100`F9d9.param`f`b`!

parameters passed to the kernel

`!`B100`F9d9.shared`f`b`!

memory shared between threads in a block

`!`B100`F9d9.tex`f`b`!

global texture memory (deprecated)

Shared memory is declared in the PTX file via lines at the start of the form:

`B100`F9d9.shared .align 8 .b8 pbatch_cache[15744]; // define 15,744 bytes, aligned to an 8-byte boundary`f`b

Writing kernels in PTX requires explicitly registering PTX modules via the CUDA Driver API, typically more cumbersome than using the CUDA Runtime API and Nvidia's CUDA compiler, nvcc. The GPU Ocelot project provided an API to register PTX modules alongside CUDA Runtime API kernel invocations, though the GPU Ocelot is no longer actively maintained.`:cite-ref-6[`F5bf`_`[6`#cite-note-6]`_`f]

>>See also

• `F33f`_`[Standard Portable Intermediate Representation`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=Standard_Portable_Intermediate_Representation]`_`f (SPIR)

• `F33f`_`[CUDA binary`:/page/wikibook/entry.mu`zim=wikipedia_en_all_nopic_2025-08.zim|entry_path=CUDA_binary]`_`f (cubin) – a type of fat binary

>>References

`:cite-note-0-1`!1.`! `F0af`_`[↑`#cite-ref-0-1-0]`_`f "User Guide for NVPTX Back-end – LLVM 7 documentation". `*llvm.org`*.

`:cite-note-2`!2.`! `F0af`_`[↑`#cite-ref-2]`_`f "CUDA Binary Utilities". `*docs.nvidia.com`*. Retrieved 2019-10-19.

`:cite-note-3`!3.`! `F0af`_`[↑`#cite-ref-3]`_`f "nvptx". `*GCC Wiki`*.

`:cite-note-4`!4.`! `F0af`_`[↑`#cite-ref-4]`_`f "Inline PTX Assembly in CUDA". `*docs.nvidia.com`*. Retrieved 2019-11-03.

`:cite-note-ptx-isa-5`!5.`! `F0af`_`[↑`#cite-ref-ptx-isa-5-0]`_`f "PTX ISA Version 2.3" (PDF).

`:cite-note-6`!6.`! `F0af`_`[↑`#cite-ref-6]`_`f "GPUOCelot: A dynamic compilation framework for PTX". `*github.com`*. 7 November 2022.

>>External links

• PTX ISA page on NVIDIA Developer Zone

`c`F0af`_`[↑ Back to top`#top]`_`f`a